//----------------------------------------------------------------------------
// M A R I T I M E  R E S E A R C H  I N S T I T U T E  N E T H E R L A N D S
//----------------------------------------------------------------------------
// Copyright (C) 2012 - MARIN - All rights reserved - http://www.marin.nl
//----------------------------------------------------------------------------
// Program    : mermaid
// Module     : vbm
// File       : boundaryzetaphi.h
// Author     : M.A. de Jong
//----------------------------------------------------------------------------

#ifndef BOUNDARYZETAPHI_H
#define BOUNDARYZETAPHI_H

#include "kernels/boundaryzetaphi_kernels.h"

namespace wavescomputer
{

namespace boundaryzetaphi
{

namespace reflectiveneumann
{

template <class T>
inline void north(T* phi, T* zeta,
           const T* phiN, const T* zetaN,
           const T* phi_in_N_ext, T* const zeta_in_N_ext,
           const T* phi_in_N_cen, T* const zeta_in_N_cen,
           const int ld,
           const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx1 / NUM_THREADS);
    
    boundaryzetaphi::reflectiveneumann::kernel_bzp_rn_north<T> 
        <<< numBlocks, NUM_THREADS >>> 
           (phi, zeta,
            phiN, zetaN, 
            phi_in_N_ext, zeta_in_N_ext,
            phi_in_N_cen, zeta_in_N_cen,
            ld, Nx1, Nx2);
    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundaryzetaphi::reflectiveneumann::kernel_bzp_rn_north");
}

template <class T>
inline void east(T* phi, T* zeta,
          const T* phiE, T* const zetaE,
          const T* phi_in_E_ext, const T* zeta_in_E_ext,
          const T* phi_in_E_cen, const T* zeta_in_E_cen,
          const int ld,
          const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx1 / NUM_THREADS);
    
    boundaryzetaphi::reflectiveneumann::kernel_bzp_rn_east<T> 
        <<< numBlocks, NUM_THREADS >>> 
           (phi, zeta,
            phiE, zetaE, 
            phi_in_E_ext, zeta_in_E_ext,
            phi_in_E_cen, zeta_in_E_cen,
            ld, Nx1, Nx2);
    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundaryzetaphi::reflectiveneumann::kernel_bzp_rn_east");
}

template <class T>
inline void south(T* phi, T* zeta,
           const T* phiS, const T* zetaS,
           const T* phi_in_S_ext, const T* zeta_in_S_ext,
           const T* phi_in_S_cen, const T* zeta_in_S_cen,
           const int ld,
           const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx1 / NUM_THREADS);
    
    boundaryzetaphi::reflectiveneumann::kernel_bzp_rn_south<T> 
        <<< numBlocks, NUM_THREADS >>> 
           (phi, zeta,
            phiS, zetaS,
            phi_in_S_ext, zeta_in_S_ext,
            phi_in_S_cen, zeta_in_S_cen,
            ld, Nx1, Nx2);
    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundaryzetaphi::reflectiveneumann::kernel_bzp_rn_south");
}

template <class T>
inline void west(T* phi, T* zeta,
          const T* phiW, const T* zetaW,
          const T* phi_in_W_ext, const T* zeta_in_W_ext,
          const T* phi_in_W_cen, const T* zeta_in_W_cen,
          const int ld,
          const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx1 / NUM_THREADS);
    
    boundaryzetaphi::reflectiveneumann::kernel_bzp_rn_west<T> 
        <<< numBlocks, NUM_THREADS >>> 
           (phi, zeta,
            phiW, zetaW, 
            phi_in_W_ext, zeta_in_W_ext,
            phi_in_W_cen, zeta_in_W_cen,
            ld, Nx1, Nx2);
    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundaryzetaphi::reflectiveneumann::kernel_bzp_rn_west");
}

} // namespace reflectiveneumann




namespace weaklyreflective
{

template <class T>
inline void north(T* phi, T* zeta,
           const T* phiN, const T* zetaN,
           const T* phi_in_N_ext, const T* zeta_in_N_ext,
           const int ld,
           const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx1 / NUM_THREADS);

    boundaryzetaphi::weaklyreflective::kernel_bzp_wr_north<T>
        <<< numBlocks, NUM_THREADS >>> 
           (phi, zeta,
            phiN, zetaN,
            phi_in_N_ext, zeta_in_N_ext,
            ld, Nx1, Nx2);
    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundaryzetaphi::weaklyreflective::kernel_bzp_wr_north");
}

template <class T>
inline void east(T* phi, T* zeta,
          const T* phiE, const T* zetaE,
          const T* phi_in_E_ext, const T* zeta_in_E_ext,
          const int ld,
          const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx1 / NUM_THREADS);

    boundaryzetaphi::weaklyreflective::kernel_bzp_wr_east<T>
        <<< numBlocks, NUM_THREADS >>> 
           (phi, zeta,
            phiE, zetaE,
            phi_in_E_ext, zeta_in_E_ext,
            ld, Nx1, Nx2);
    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundaryzetaphi::weaklyreflective::kernel_bzp_wr_east");
}

template <class T>
inline void south(T* phi, T* zeta,
           const T* phiS, const T* zetaS,
           const T* phi_in_S_ext, const T* zeta_in_S_ext,
           const int ld,
           const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx1 / NUM_THREADS);

    boundaryzetaphi::weaklyreflective::kernel_bzp_wr_south<T>
        <<< numBlocks, NUM_THREADS >>> 
           (phi, zeta,
            phiS, zetaS,
            phi_in_S_ext, zeta_in_S_ext,
            ld, Nx1, Nx2);
    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundaryzetaphi::weaklyreflective::kernel_bzp_wr_south");
}

template <class T>
inline void west(T* phi, T* zeta,
          const T* phiW, const T* zetaW,
          const T* phi_in_W_ext, const T* zeta_in_W_ext,
          const int ld,
          const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx1 / NUM_THREADS);

    boundaryzetaphi::weaklyreflective::kernel_bzp_wr_west<T>
        <<< numBlocks, NUM_THREADS >>> 
           (phi, zeta,
            phiW, zetaW,
            phi_in_W_ext, zeta_in_W_ext,
            ld, Nx1, Nx2);
    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundaryzetaphi::weaklyreflective::kernel_bzp_wr_west");
}

} // namespace weaklyreflective

} // namespace boundaryzetaphi

} // namespace wavescomputer

#endif // BOUNDARYZETAPHI_H

